#ifndef PACK_CL
#define PACK_CL
#include <types.cl>
#include <constants.cl>

__kernel void pack(__global unsigned int* mz,
                   __global unsigned int* peaks,
                   unsigned int scanStride,
                   __global MSScanInfo* scanInfo,
#if defined(GLOBAL_MEMORY_SPECTRUM) || defined(LOCAL_MEMORY_SPECTRUM)
                   __global unsigned int* mzOut,
                   __global unsigned int* peaksOut
#else
                   write_only image2d_t mzOut,
                   write_only image2d_t peaksOut
#endif
                   )
{
  unsigned int mySize = scanInfo[get_global_id(1)].mNumPeaks;
  unsigned int myScan = get_global_id(1);

  //Don't copy MS 1s because we don't need them
  if(scanInfo[get_global_id(1)].mMSLevel == 1)
  {
    return;
  }

#if defined(GLOBAL_MEMORY_SPECTRUM) || defined(LOCAL_MEMORY_SPECTRUM)
  //Convert from numScans x scanStride array to numScans x kMaxPeakCount array
  for(unsigned int curPeak = get_local_id(0); 
      curPeak < mySize; 
      curPeak += get_local_size(0))
  {
    mzOut[kMaxPeakCount * myScan + curPeak] = mz[scanStride * myScan + curPeak];
    
    peaksOut[kMaxPeakCount * myScan + curPeak] = 
      peaks[scanStride * myScan + curPeak];
  }

#elif defined(IMAGE_SPECTRUM)
  //Write spectrum data into images
  int2 idx = {0, get_global_id(1)};
  uint4 val = {0.0f, 0.0f, 0.0f, 0.0f};
  
  mz = &mz[scanStride * idx.y];
  peaks = &peaks[scanStride * idx.y];
  

  for(idx.x = get_local_id(0); 
      idx.x < mySize; 
      idx.x += get_local_size(0))
  {
    val.x = mz[idx.x];
  
    write_imageui(mzOut, idx, val);

    val.x = peaks[idx.x];

    write_imageui(peaksOut, idx, val);
  }
#else
#error Unknown spectrum packing method
#endif
}

#endif

